__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;

__kernel void test_kernel(__global const float *a, __global const float *b, __global float *result)
{
    int gid = get_global_id(0);

    result[gid] = a[gid] + b[gid];
}

__kernel void ElementAdd(__read_only image2d_t input_a, __read_only image2d_t input_b, __write_only image2d_t output,
                         const int2 output_shape, float act_min, float act_max)
{
    int X = get_global_id(0) - 1;
    int Y = get_global_id(1) - 1;
    if (X >= output_shape.x || Y >= output_shape.y)
    {
        return;
    }

    float4 a = read_imagef(input_a, smp_none, (int2)(X, Y));
    float4 b = read_imagef(input_b, smp_none, (int2)(X, Y));
    float4 result = a - b;
    result = clamp(result, (float)(act_min), (float)(act_max));
    write_imagef(output, (int2)(X, Y), result);
}
